#include "glWorld.cuh"
#include <stdio.h>

#define CREATE_AND_LOAD_MATRIX(m, t, i)                    \
	matrix4x4 m;                                             \
	                                                         \
	m.row[0] = tex1Dfetch(t, i+0);                           \
	m.row[1] = tex1Dfetch(t, i+1);                           \
	m.row[2] = tex1Dfetch(t, i+2);                           \
	m.row[3] = tex1Dfetch(t, i+3);                           \

texture<float4, 1, cudaReadModeElementType> g_texMatrix;
texture<float4, 1, cudaReadModeElementType> g_texPosition;

__global__ 
void Kernel_Pos_TransformWS(float4 * pPosW, 
													  float4 * pPosR, 
														uint32 * pPosID, 
														uint32   nPos)
{
	uint32 tid = blockIdx.x * blockDim.x + threadIdx.x;

	if (tid < nPos)
	{
		uint32 matrixid = pPosID[tid];

		CREATE_AND_LOAD_MATRIX(mWS, g_texMatrix, matrixid*4);	

		pPosW[tid] = pPosR[tid] * mWS;
	}
}

__global__ 
void Kernel_Dir_TransformWS(float4 * pDirW, 
														float4 * pDirR, 
														uint32 * pPosID, 
														uint32   nPos)
{
	uint32 tid = blockIdx.x * blockDim.x + threadIdx.x;

	if (tid < nPos)
	{
		uint32 matrixid = pPosID[tid];

		CREATE_AND_LOAD_MATRIX(mWS, g_texMatrix, matrixid);	

		float4 dir;
		
		dir = normalize4(Matrix4x4TransformDirectionL(mWS, pDirR[tid]));

		pDirW[tid] = dir;
	}
}

void Wrapper_TransformWS(uint32   nPos, 
												 uint32   nMatrix,
												 uint32 * pPosID,
												 float4 * pMatrix,
												 float4 * pPosW, float4 * pPosR, 												  
												 float4 * pNmlW, float4 * pNmlR,
												 float4 * pTntW, float4 * pTntR,
												 float4 * pBntW, float4 * pBntR)
{
	if (pPosR != NULL && pPosW != NULL)
	{	
		dim3 threads(128, 1, 1);

		dim3 blocks((nPos + threads.x - 1) / threads.x, 1);

		uint32 offset = 0;

		cudaBindTexture(&offset, g_texMatrix, pMatrix, cudaCreateChannelDesc<float4>(), nMatrix*sizeof(matrix4x4));

		Kernel_Pos_TransformWS<<<blocks,threads>>>(pPosW, pPosR, pPosID, nPos);

		if (pNmlR != NULL && pNmlW != NULL)
		Kernel_Dir_TransformWS<<<blocks,threads>>>(pNmlW, pNmlR, pPosID, nPos);

		if (pTntR != NULL && pTntW != NULL)
		Kernel_Dir_TransformWS<<<blocks,threads>>>(pTntW, pTntR, pPosID, nPos);

		if (pBntR != NULL && pBntW != NULL)
		Kernel_Dir_TransformWS<<<blocks,threads>>>(pBntW, pBntR, pPosID, nPos);

		cudaThreadSynchronize();
	}
}

__global__
void Kernel_PrimitiveAABB(uint32   nPrim, 
													uint32 * pPosIdxSwzld,
													float2 * pSplit_X_LT,
													float2 * pSplit_Y_LT,
													float2 * pSplit_Z_LT)
{
	uint32 tid = blockIdx.x * blockDim.x + threadIdx.x;

	if (tid < nPrim)
	{
		float4 p0 = tex1Dfetch(g_texPosition, pPosIdxSwzld[tid+nPrim*0]);
		float4 p1 = tex1Dfetch(g_texPosition, pPosIdxSwzld[tid+nPrim*1]);

		float3 pmin = min3(p0, p1);
		float3 pmax = max3(p0, p1);

		float4 p2 = tex1Dfetch(g_texPosition, pPosIdxSwzld[tid+nPrim*2]);

		pmin = min3(pmin, p2);
		pmax = max3(pmax, p2);

		pSplit_X_LT[tid] = f2(pmin.x, tid);
		pSplit_Y_LT[tid] = f2(pmin.y, tid);
		pSplit_Z_LT[tid] = f2(pmin.z, tid);

		pSplit_X_LT[tid + nPrim] = f2(pmin.x, tid);
		pSplit_Y_LT[tid + nPrim] = f2(pmin.y, tid);
		pSplit_Z_LT[tid + nPrim] = f2(pmin.z, tid);

	}
}

void Wrapper_PrimitiveAABB(uint32   nPrim, 
													 uint32   nPos,
													 float4 * pPos,
													 uint32 * pPosIdxSwzld,
													 float2 * pSplit_X,
													 float2 * pSplit_Y,
													 float2 * pSplit_Z)
{
	dim3 threads(128, 1, 1);
	dim3 blocks( (nPrim + threads.x - 1 / threads.x), 1);

	cudaBindTexture(NULL, g_texPosition, pPos, cudaCreateChannelDesc<float4>(), nPos * sizeof(float4));

	Kernel_PrimitiveAABB<<<blocks,threads>>>(nPrim, 
		                                       pPosIdxSwzld, 
																					 pSplit_X, 
																					 pSplit_Y, 					
																					 pSplit_Z);
	cudaThreadSynchronize();
}